已經 很多天的 兩萬字數AI slob了
之所以說是 AI slob，是因為這幾篇文字真的很多、很滿。

其實我覺得自己講了太多數學和結構層面，反而沒有著墨太多硬體優化與架構優化的部分，這是我這 10 天以來比較疏忽的地方。我原本努力想在每一天的每個 section 去說明今天教了什麼、GPU 在講什麼、哪裡做了優化，但後來發現內容有點老生常談、重複，中間還夾雜了一堆跟 GPU 加速無關的數學觀念，換句話說，有點像廢文。不過自從聽了某位大佬說「反正鐵人賽不就是寫一些有趣的廢文給大家廢文共賞嗎，幹嘛那麼緊張」，我也就釋懷了。

我對自己的期許其實很單純：

1. 至少發出來的東西是我親手做過、大概知道在幹嘛的。(親手跟AI一起做 哈哈)
2. 雖然不敢說 100% 掌握每個細節、關鍵字的因果關係或 branch test，但至少文章裡提到的因果邏輯，我是能講出一兩分道理的。

這 9 天實際上講了三個案例，但因為是同一個主題，所以也可以算是一個大案例：
• 前兩天是小型 kernel
• 後來是一個大型 kernel
• 最後是把一整個完整演算法搬到 GPU 上重作

我必須承認，這幾天寫下來真的很像大學生自幹的 side hustle，
當然啦 實際上也就是研究生自幹的 side hustle，沒什麼太大差異。

明天就是第 30 天了，更多的 feedback 和總結感想我留到明天再寫。時間也快到了（大概就是今天晚上 12 點吧），我會早早寫完發出去。另外，這幾篇 title 我都標註「重賽版」，希望之後有時間能靜下心來，把每一篇重新整理成更完整的好文章，到時就不再只是重賽版，而是完整版概念的內容。

來吧，我們開始第九天，把這個 title 案子、這整個 LCN 的總結在今天 close 掉。我們開始吧！

前面八篇依序做了三類改變：

Crossing
    full recount
        -> incremental update

Layout
    repeatedly load every position from global memory
        -> register accumulation + shared-memory tiling

Simulated Annealing
    one host decision per proposal
        -> device-resident loop + tempered replicas

這些改變都可能讓程式變快，也都可能讓程式更快地算出錯誤答案。

今天只處理一件事：一個 GPU optimization 完成後，怎麼提出可信的效能結論？

我們會用實際保存的 crossing benchmark 回答：

48 nodes / 96 edges
    full kernel  = 0.02473 ms
    delta kernel = 0.02243 ms
    speedup      = 1.10x

128 nodes / 512 edges
    full kernel  = 0.25788 ms
    delta kernel = 0.13190 ms
    speedup      = 1.96x

第二組數字是真的，但它只回答一個非常窄的問題：

在 RTX 4060 Laptop GPU 上，
固定同一個 base layout 與 32 個 independent proposals，
只計 CUDA kernel execution，
incremental crossing evaluator 相對 full evaluator 快多少？

它沒有測量完整 LCN solver。

效能問題其實有三層

看到一個 speedup 之前，先問它回答哪一層。

Layer 1: evaluator correctness
    CPU 與 GPU 是否算出相同的 per-edge crossing counts？

Layer 2: isolated mechanism
    incremental、tiling 或 fusion 是否讓指定 kernel 更快？

Layer 3: solver outcome
    在相同時間內，是否更常、更早找到更好的合法 layout？

三層需要不同的輸入、計時器和判定方式。

把 Layer 2 的答案直接當成 Layer 3，是 GPU benchmark 最常見的誤導。

為什麼選 Crossing Evaluator 當主案例？

Day 29 需要一個工作量能完全鎖住的例子。

Crossing evaluation 很適合，因為兩個 kernels 可以接收完全相同的輸入，並產生完全相同形狀的輸出：

input
    one immutable base layout
    one graph
    32 node-move proposals

full evaluator output
    32 * E per-edge counts

delta evaluator output
    32 * E per-edge counts

兩條路只差計算方法：

full
    recompute every edge against every other edge

delta
    reuse base counts
    recompute pairs affected by the moved node

這讓我們能先比整個輸出陣列，再量同一份工作的 kernel time。

Layout 與 SA 就沒有這麼單純。Force-directed layout 的 floating-point reduction order 可能改變軌跡；SA 更會因 RNG 與接受歷史不同而走到不同狀態。它們最後必須用品質分布與 time-to-target 比較。

先從 deterministic evaluator 建立 benchmark discipline，再處理 stochastic solver，是這個例子的選擇理由。

第一個 Gate：先證明算的是同一件事

Crossing demo 同時保留四條計算路徑：

CPU full recount
CPU pair-based delta update
GPU full recount
GPU incident-edge delta update

對每一個 proposal，要求：

CPU_full_counts[proposal][edge]
    == CPU_delta_counts[proposal][edge]
    == GPU_full_counts[proposal][edge]
    == GPU_delta_counts[proposal][edge]

比較單位是每條 edge，不只是最大值 K。

假設正確答案是：

counts = [3, 3, 1, 1]
K      = 3

錯誤 kernel 回傳：

counts = [3, 2, 2, 1]
K      = 3

只比 K 會通過，但 n_K、Phi、C 全部已經不同：

correct objective = (3, 2, 20, 4)
wrong objective   = (3, 1, 18, 4)

SA 會依錯誤 objective 做出不同的 acceptance decision。錯誤不一定立刻出現在最終 K，卻已經改變搜尋路徑。

CPU Oracle 必須和 GPU 共享定義，不共享 Bug

crossing_demo.py 的 CPU oracle 使用 Python arbitrary-precision integers 計算 orientation：

def orient(p, q, r):
    return ((q[0] - p[0]) * (r[1] - p[1])
            - (q[1] - p[1]) * (r[0] - p[0]))

教學 CUDA kernel 則先轉成 long long：

__device__ long long orient(
    int ax, int ay,
    int bx, int by,
    int cx, int cy
) {
    return ((long long)bx - ax) * ((long long)cy - ay)
         - ((long long)by - ay) * ((long long)cx - ax);
}

如果 GPU 版本先用 32-bit integer 完成乘法，再把結果存入 64-bit，座標較大時仍可能先 overflow：

wrong
    long long value = int_difference * int_difference

correct direction
    cast before multiplication

Oracle 和 optimized path 可以實作同一個數學定義，但不應直接呼叫同一個低階函式。否則一個 shared bug 會讓 differential test 看起來全綠。

幾何 Fixtures 要覆蓋 Boundary

隨機圖適合壓力測試，卻不保證產生每種幾何關係。Demo 先固定小型 fixtures：

proper interior crossing
shared endpoint
collinear contact
positive-length collinear overlap
isolated vertex
no-op move
move that removes current maximum
move that creates a new crossing

這些 fixtures 確認 crossing 定義，再用較大 random workload 做 differential comparison。

但這個教學 evaluator 只判斷 strict proper crossing。它不是完整 LCN validator。

完整 solver 發布答案之前還要獨立檢查：

V1  every node is inside the allowed grid
V2  no duplicate node coordinates
V3  no non-endpoint node lies on an edge
V4  no two edges overlap along positive length

因此 correctness gates 是分層的：

per-edge crossing parity
    -> objective parity
        -> candidate geometry validity
            -> final published layout revalidation

前一層通過，不能省略下一層。

為什麼不能直接用 CPU Timer 包住 CUDA Launch？

CUDA kernel launch 預設是 asynchronous。

下面的寫法主要量到 enqueue：

started = time.perf_counter()
kernel(grid, block, args)
elapsed = time.perf_counter() - started

CPU 在 GPU 完成之前就可能繼續執行，所以這個 elapsed 不是 kernel duration。

如果要量 end-to-end host-visible latency，可以在結尾同步：

started = time.perf_counter()
kernel(grid, block, args)
stream.synchronize()
elapsed = time.perf_counter() - started

如果要量 device execution，使用 CUDA events：

start = cp.cuda.Event()
stop = cp.cuda.Event()

start.record()
for _ in range(100):
    kernel(grid, block, args)
stop.record()
stop.synchronize()

milliseconds_per_launch = (
    cp.cuda.get_elapsed_time(start, stop) / 100
)

兩種時間都可以報，但回答不同問題。

三種 Timing Scope 要分開命名

Kernel-only

CUDA event starts
    -> repeated kernel launches
CUDA event stops

通常排除：

Python startup
CUDA context creation
CuPy JIT compilation
cudaMalloc
H2D input transfer
D2H output transfer
CPU validation
file output

它最適合回答某個 kernel mapping 是否改善。

GPU stage

prepare already-loaded device state
    -> launch one or more kernels
    -> synchronize
    -> expose stage result

它可能包含 launch gaps、exchange boundaries、stream synchronization 和必要的 device copies。

End-to-end solver

read instance
    -> generate initial layouts
    -> validate and admit seeds
    -> allocate/upload
    -> run search
    -> repair
    -> final exact validation
    -> save output

競賽或產品使用者真正等待的是這一層。

這三個數字應該像不同單位一樣清楚標示，不能拿 CPU end-to-end 除以 GPU kernel-only 得到 speedup。

Warm-up 不是為了讓數字更好看

第一次 GPU 呼叫可能包含：

CUDA context initialization
module loading
JIT compilation
memory pool growth
page mapping
clock ramp-up
cache cold start

如果研究 steady-state kernel throughput，應先 warm up，再開始 CUDA events。

這次 crossing benchmark 使用：

5 warm-up launches
9 timing batches
100 launches per batch
report median batch time per launch

為什麼每個 sample 包 100 launches？

小 kernel 只有數十 microseconds。把多次 launches 放進同一對 events，可以降低 event resolution 與單次排程抖動在比例上的影響。

為什麼報 median？

某次 batch 可能受到 OS scheduling、背景 GPU 工作或動態時脈影響。Median 比 mean 不容易被單一長尾拖動。

不過 warm-up 的使用要和問題一致。

question: steady-state kernel throughput
    exclude one-time JIT and initialization

question: command-line tool first-run latency
    include them

沒有一種 scope 永遠正確，只有 scope 是否和結論一致。

實測環境與結果

保存的環境資料是：

GPU          NVIDIA GeForce RTX 4060 Laptop GPU
CUDA runtime 12.9
CuPy         13.6.0
Python       3.11.15 under WSL
seed         42
proposals    32 independent moves

兩組 CUDA-event 結果：

兩組都先通過：

CPU full
    == CPU delta
    == GPU full
    == GPU delta

這組數字展示一個重要的 GPU 現象。

小 workload 時：

saved pair tests are limited
fixed launch and control cost remains

-> only 1.10x

較大 workload 時：

full evaluator grows with all edge pairs
delta evaluator reuses base counts
affected work is tied to incident edges

-> 1.96x in this measurement

演算法減少的工作量必須大到足以超過固定成本，speedup 才會明顯。這也是為什麼不能只用 big-O 宣稱實際 GPU 倍率。

這個 1.96x 可以支持哪些句子？

可以說：

On the saved 128-node, 512-edge teaching workload,
the incremental proper-crossing CUDA kernel had a
1.96x lower median kernel time than the full kernel.

不能說：

the complete LCN solver is 1.96x faster

也不能說：

incremental crossing is always 1.96x faster

更不能說：

the optimized solver finds solutions 1.96x better

實際報告已把 scope 寫成：

single-layer proper-crossing evaluator;
independent proposals;
not official solver

Timing scope 也明確寫出：

CUDA events;
5 warmups;
9 batches x 100 launches;
excludes allocation/JIT/transfers/CPU verification

Benchmark 的限制不是附註。它是結論本身的一部分。

新增一個會拒絕錯誤報告的程式

今天新增：

• benchmark_report.py

它不重新執行 GPU benchmark，而是稽核保存的 JSON。只有以下條件全部通過，才輸出 speedup：

CPU full equals CPU delta
GPU full equals CPU
GPU delta equals CPU
all timing values are positive and finite
min <= median <= max
recorded speedup equals full_median / delta_median
timing scope explicitly says CUDA events
all aggregated reports use the same device and scopes

核心 gate 是：

failed = [
    key for key in checks
    if report[key] is not True
]

if failed:
    raise ReportError(
        "correctness gate failed: " + ", ".join(failed)
    )

Speedup 也從兩個 median 重新計算：

recomputed = full_ms / delta_ms

if not math.isclose(recomputed, recorded):
    raise ReportError("recorded speedup does not match medians")

這避免手動更新表格時把舊的 speedup 留在新數據旁邊。

執行 Report Audit

在 repository root 執行：

python case_4/examples/benchmark_report.py

輸出：

| workload | GPU full median | GPU delta median | full / delta |
|---|---:|---:|---:|
| 48 nodes / 96 edges / 32 proposals | 0.02473 ms | 0.02243 ms | 1.10x |
| 128 nodes / 512 edges / 32 proposals | 0.25788 ms | 0.13190 ms | 1.96x |

Correctness gate: PASS (CPU full == CPU delta == GPU full == GPU delta)
Device: NVIDIA GeForce RTX 4060 Laptop GPU
Evaluator scope: single-layer proper-crossing evaluator; independent proposals; not official solver
Timing scope: CUDA events; 5 warmups; 9 batches x 100 launches; excludes allocation/JIT/transfers/CPU verification
Claim boundary: kernel-only evaluator result; not full-solver speedup.

若任一 correctness flag 是 false，程式會回傳 non-zero exit code，也不產生 aggregate claim：

REJECT report.json: correctness gate failed: gpu_delta_equals_cpu
No aggregate claim: at least one report failed the benchmark contract.

這個例子介紹的不是 CUDA arithmetic，而是效能資料的 admission control。錯誤結果不應進入報表，更不應進入 optimization decision。

現有報告還不能回答統計顯著性

兩份 JSON 保存：

median
minimum
maximum

但沒有保存 9 個原始 batch samples。

因此可以重算 median-based speedup，卻不能從目前檔案可靠地建立 bootstrap confidence interval，也不能檢查完整分布是否雙峰或持續 drift。

更完整的 benchmark report 應保存：

{
  "warmups": 5,
  "launches_per_batch": 100,
  "full_samples_ms": [0.25, 0.26, 0.25],
  "delta_samples_ms": [0.13, 0.13, 0.14]
}

最好還要交錯執行 A/B：

full, delta, delta, full, ...

或隨機化順序，降低 GPU 溫度、時脈和背景負載隨時間變化造成的偏差。

目前資料足以做窄範圍的 teaching comparison；若要決定 production default，應重新執行並保存 raw samples。

Laptop GPU 還有一個容易忽略的變因

這次裝置是 Laptop GPU。相同 kernel 可能因以下因素改變：

AC power or battery
performance mode
GPU temperature
dynamic boost clock
CPU/GPU shared cooling
another process using the GPU
display workload

因此環境記錄至少要包含：

GPU model
driver and CUDA runtime
library version
power mode
benchmark timestamp
source commit or file hash
binary hash
command line
input hash
seed list

只有 GPU 名稱仍不足以保證兩次 run 完全可比，但比完全沒有 provenance 好得多。

Throughput 也必須有明確的 Work Unit

常見報告會寫：

1,000,000 iterations / second

但一個 iteration 可能是：

one proposal generated
one legal proposal evaluated
one accepted transition
one replica step
one group round containing 32 candidates

這些不是同一個 work unit。

Day 27 的 parallel proposal group 每輪可能評估 32 個 candidates，最後只 commit 一個 state。若拿 group rounds/sec 和單鏈 proposals/sec 比較，數字沒有意義。

建議同時報：

attempted proposals / second
valid evaluated proposals / second
accepted moves / second
dependent chain steps / second

對 PT 再加：

replica-local steps / second
exchange attempts / second
round trips / wall-clock budget

Throughput 說明硬體處理多少工作；它本身不說明這些工作對解題有沒有價值。

固定 Proposal Count 還是固定 Wall-clock？

兩種 benchmark 都需要，但回答不同問題。

固定工作量

same base layout
same proposals
same output shape
same correctness oracle

適合比較：

full crossing vs delta crossing
naive force load vs tiled force load
unfused vs fused kernel sequence

固定時間

same instance
same starting portfolio policy
same wall-clock budget
multiple paired seeds

適合比較：

single SA vs resident SA
independent replicas vs parallel tempering
operator policy A vs operator policy B

固定 proposals 對 stochastic search 仍有分析價值，但它可能懲罰每一步較貴、卻更有方向性的 operator。固定 wall-clock 才回答真實部署時能交出什麼答案。

Stochastic Solver 不能只跑一次

假設兩個版本各跑一個 seed：

baseline best K  = 4
optimized best K = 3

不能立刻說 optimized 比較好。這可能只是 RNG 差異。

應使用 paired seed list：

seed 0x2901
seed 0x2902
seed 0x2903
...

每個 candidate 都跑相同 seeds，並交錯或隨機安排 candidate 順序。

對每個 run 保存：

initial validated objective
final validated objective
time to first target K
time to best K
best stage
actual wall time
exit status
validation errors

最後比較分布：

success rate within budget
median final K
median full objective among equal K
median time-to-target among successes
all failures and timeouts retained

不要刪除沒達標的 seeds。刪除 failures 會讓最不穩定的演算法看起來最快。

Time-to-target 比 Final K 更能看出搜尋速度

假設三個版本在 30 秒後都得到 K = 3：

baseline
    first K=3 at 27.2 s

resident SA
    first K=3 at 9.4 s

parallel tempering
    first K=3 at 3.1 s

只看 final K，三者平手。Time-to-target 才顯示 anytime behavior。

但 target 必須事先定義：

target K = 3

不能看完結果後再挑最有利的 threshold。

對未達標的 run，要保留 censored 狀態：

time_to_target = NOT_REACHED_WITHIN_30S

不能把它填成 0，也不能從 median 中默默移除。

Best K 相同，不代表答案相同

LCN 正式 objective 是：

(K, n_K, Phi, C)

兩個 versions 都回報 K = 3 時，仍可能是：

version A = (3, 12, 420, 81)
version B = (3,  4, 180, 40)

Version B 在正式 lexicographic ranking 中更好。

因此 solver report 至少要保存完整 validated tuple。只保存 best_k 會丟失同一 K 內的 progress，也無法判斷 cold replicas 是否真的在精修。

此外，kernel 自己回報的 best 不是最後真相：

kernel_best
    candidate claimed by optimized device path

verified_best
    coordinates re-evaluated by official checker

兩者不一致時應拒絕 publication，保留 diagnostic，而不是相信比較小的數字。

1.96x 的 Kernel Speedup 會怎麼影響完整 Solver？

假設 crossing kernel 只佔完整執行時間的百分之四十，其餘百分之六十完全沒變。即使 crossing 加速 1.96x，整體理論速度也只有：

fraction_optimized = 0.40
kernel_speedup     = 1.96

total_speedup
    = 1 / ((1 - fraction_optimized)
           + fraction_optimized / kernel_speedup)

    = 1 / (0.60 + 0.40 / 1.96)
    = 1.24x

如果 crossing 原本只佔百分之十：

total_speedup
    = 1 / (0.90 + 0.10 / 1.96)
    = 1.05x

這就是為什麼要先 profile complete pipeline。最佳化一個看起來很核心、實際只佔少量時間的 kernel，端到端結果可能幾乎不變。

反過來，Day 27 的 resident loop 可能沒有大幅縮短單一 crossing kernel，卻因為移除成千上萬次 host synchronization 而明顯改善 pipeline time。

NCU 應該用來驗證哪個假設？

Correctness 與 wall-clock 結果確認後，再用 Nsight Compute 找原因。

Day 23：Incremental Crossing

要驗證：

edge-pair tests decrease
global reads tied to unaffected edges decrease
branch divergence does not erase the saved work

Day 25：Shared-memory Tiling

要驗證：

global position-load requests decrease
shared-memory traffic increases as intended
barriers and bank conflicts remain controlled
register use does not collapse occupancy

Day 27：Resident SA

要驗證：

kernel launch count decreases
CPU/GPU gaps decrease
timed allocations are zero
state remains in device memory across decisions

Day 28：Parallel Tempering

要驗證：

replica blocks keep the GPU occupied
exchange kernel traffic matches accepted swaps
workspace fits without unexpected spill or allocation churn
extra throughput produces useful round trips

不要先看到 occupancy 低就盲目調 block size。Occupancy 是限制線索，不是最終 KPI。Kernel 可能在較低 occupancy 下已有足夠 latency hiding，也可能真正受限於 memory bandwidth、integer arithmetic 或 synchronization。

Counterintuitive：Kernel 變慢，Solver 也可能變好

假設 directional proposal 比 uniform proposal 每一步多做一些計算：

uniform
    2,000,000 proposals / second
    target success rate = 10%

directional
    1,200,000 proposals / second
    target success rate = 70%

Directional kernel throughput 較低，卻可能更快找到好 layout。

另一個例子是 PT exchange：交換完整 state 會增加 global-memory traffic，單看 local steps/sec 可能下降；如果 round trips 幫助 cold replica 更早找到 K = 3，solver 仍然更有效。

所以 GPU search optimization 必須同時保留兩類指標：

mechanism metrics
    kernel time
    launches
    bytes
    proposals/sec

outcome metrics
    validated objective
    success rate
    time-to-target

前者解釋硬體行為，後者決定最佳化是否值得。

反過來：Throughput 變高，Solver 也可能變差

如果為了 throughput 改變 proposal semantics：

before
    focus on bottleneck edge endpoints

after
    choose any node uniformly because it is cheaper

每秒 proposal 數可能上升，但大部分 moves 和目前的 worst edges 無關。

如果為了避免 validation 分支，直接跳過 V3/V4 checks，數字會更漂亮，輸出卻可能不合法。

如果把 512 proposal workers 誤當 512 private chains，報告的「parallelism」也會高估搜尋 breadth。

這些都是同一個問題：work unit 或 semantics 改變後，舊 denominator 已經不能使用。

一份可審查的 Benchmark Manifest

完整實驗應保存 machine-readable manifest：

{
  "source_commit": "...",
  "dirty_worktree": false,
  "binary_sha256": "...",
  "dataset_ids": ["instance_01", "instance_06"],
  "input_sha256": {"instance_01": "..."},
  "candidate_configs": {
    "baseline": {"incremental": false},
    "candidate": {"incremental": true}
  },
  "seed_list": [10497, 10498, 10499],
  "run_order": ["baseline", "candidate"],
  "wall_budget_seconds": 30,
  "raw_run_paths": ["raw/baseline_10497.json"]
}

Manifest 解決的是「這個數字從哪裡來」。Summary table 不能取代 raw runs，README 也不能取代 binary hash。

專案現有 HO-037 harness 已經採用這個方向：raw run、manifest、summary 與 promotion decision 分開保存。如果 evaluator matrix 尚未實作，它會明確輸出：

INSUFFICIENT_EVIDENCE

這比用缺少的資料猜一個結論可靠。

最後的 Benchmark Checklist

在計時之前

write the mathematical contract
build an independent oracle
cover geometry boundary fixtures
compare full per-edge outputs
run final V1-V4 validation

對 deterministic kernels

same input arrays
same proposal list
same output shape
warm up deliberately
use CUDA events
batch very small launches
save raw timing samples

對 stochastic solvers

same instance set
paired seed list
same wall-clock budget
randomized or interleaved run order
retain failures and timeouts
report full validated objective
report time-to-target

在發布 speedup 時

name numerator and denominator
state timing scope
state excluded work
record hardware and software environment
link raw reports
limit the claim to the measured workload

這九篇真正完成的路線

整個 LCN GPU case 並不是先寫 CUDA，再找一個數字證明它快。

Day 21
    define graph, edge, legal layout, crossing and LCN objective

Day 22-23
    build a full crossing oracle
    then reduce repeated edge-pair work incrementally

Day 24-25
    derive force-directed layout
    then tile reused coordinates through shared memory

Day 26-28
    define one SA chain
    keep its control on device
    then add private replicas and temperature exchange

Day 29
    prove output parity
    separate timing scopes
    measure search quality under equal budgets

最值得保留的工作順序是：

define the exact answer
    -> write a readable baseline
        -> identify repeated computation and movement
            -> map reuse and ownership onto GPU memory
                -> prove semantics did not change
                    -> measure the layer actually changed
                        -> measure end-to-end outcome

GPU optimization 不只是把 arithmetic 分給更多 threads。它同時在安排：

who owns each state
where data stays
how often data moves
which work can run concurrently
where synchronization is unavoidable
what result is allowed to be published

只有 correctness、measurement scope 與解題品質都能對上，1.96x 才是一個工程結論，而不只是一個漂亮數字。

程式與原始資料

• Day 29 benchmark report audit
• CPU/GPU differential benchmark
• 教學 CUDA full/delta kernels
• 48 nodes / 96 edges 原始報告
• 128 nodes / 512 edges 原始報告
• 保存的執行環境
• 完整 solution validator
• Benchmark manifest contract tests
• Replica 與 exchange telemetry contract tests